#include <cutil_inline.h>
#include <cutil_math.h>


texture<uchar4, 2, cudaReadModeNormalizedFloat>	rgbaTex;
cudaArray*	d_array;


__device__ uint RGBAFloatToUint(float4 rgba)
{
	rgba.x = __saturatef(rgba.x);
	rgba.y = __saturatef(rgba.y);
	rgba.z = __saturatef(rgba.z);
	rgba.w = __saturatef(rgba.w);
	return (uint(rgba.w * 255.0f) << 24) | (uint(rgba.z * 255.0f) << 16) | (uint(rgba.y * 255.0f) << 8) | uint(rgba.x * 255.0f);
}

__device__ float4 RGBAUintToFloat(uint c)
{
	float4 rgba;
	rgba.x = (c & 0xff) * 0.003921568627f;       //  /255.0f;
	rgba.y = ((c>>8) & 0xff) * 0.003921568627f;  //  /255.0f;
	rgba.z = ((c>>16) & 0xff) * 0.003921568627f; //  /255.0f;
	rgba.w = ((c>>24) & 0xff) * 0.003921568627f; //  /255.0f;
	return rgba;
}


extern "C" void InitTexture(const void* imageData, int width, int height)
{
	const int size = width * height * sizeof(unsigned int);

	// copy image data to array
	cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsigned);
	cutilSafeCall(cudaMallocArray(&d_array, &channelDesc, width, height)); 
	cutilSafeCall(cudaMemcpyToArray(d_array, 0, 0, imageData, size, cudaMemcpyHostToDevice));

	// Bind the array to the texture
	cutilSafeCall(cudaBindTextureToArray(rgbaTex, d_array));
}
extern "C" void FreeTexture()
{
	cutilSafeCall(cudaFreeArray(d_array));
}


__device__ float4 MakeGray(float4 rgba)
{
	float lum = rgba.x * 0.3f + rgba.y * 0.59f + rgba.z * 0.11f;
	return make_float4(lum, lum, lum, rgba.w);
}

__device__ float4 MakeYUV(float4 rgba)
{
	const float Kr = 0.299f;
	const float Kb = 0.114f;
	const float Kg = 1.0f - Kb - Kr;

	const float Y = Kr * rgba.x + Kg * rgba.y + Kb * rgba.z;
	const float U = 0.436f * ((rgba.y - Y) / (1.0f - Kb));
	const float V = 0.615f * ((rgba.x - Y) / (1.0f - Kr));

	return make_float4(Y, U, V, rgba.w);
}

__global__ void kernel_ProcessImage(uint* outData, int width, int height)
{
	const uint x = __umul24(blockIdx.x, blockDim.x) + threadIdx.x;
	const uint y = __umul24(blockIdx.y, blockDim.y) + threadIdx.y;

	if (x < width && y < height)
	{
		float4 pixel = tex2D(rgbaTex, x, y);
		outData[y * width + x] = RGBAFloatToUint(MakeYUV(pixel));
	}
}

extern "C" void ProcessImage(uint* destData, int width, int height)
{
	dim3 gridSize(width / 16, height / 16);
	dim3 blockSize(16, 16);
	kernel_ProcessImage<<<gridSize, blockSize>>>(destData, width, height);
}
